

// 9. source vector is accesed as uint4 vector

__kernel void minp( __global uint4 *src, 
	 __global uint *gmin, 
	 __local uint *lmin, 
	 uint nitems ) 
{ 
 // 10. Set up __global memory access pattern. 
 
 uint global_size = get_global_size(0);
 uint count = ( nitems / 4 ) / global_size; 
 uint idx =  get_global_id(0); 
 uint stride =  global_size; 
 uint pmin = (uint) -1; 
 
 // 11. First, compute private min, for this work-item. 
 
 for( int n=0; n < count; n++, idx += stride ) 
 { 
  pmin = min( pmin, src[idx].x ); 
  pmin = min( pmin, src[idx].y ); 
  pmin = min( pmin, src[idx].z ); 
  pmin = min( pmin, src[idx].w ); 
 } 
 
 // 12. Reduce min values inside work-group. 
 
 if( get_local_id(0) == 0 ) 
  lmin[0] = (uint) -1; 
 
 barrier( CLK_LOCAL_MEM_FENCE ); 
 
 (void) atom_min( lmin, pmin ); 
 
 barrier( CLK_LOCAL_MEM_FENCE ); 
 
 // Write out to __global. 
 
 if( get_local_id(0) == 0 ) 
  gmin[ get_group_id(0) ] = lmin[0]; 

}


